package hpc.exercices.exercice2;

import static org.jocl.CL.CL_CONTEXT_PLATFORM;
import static org.jocl.CL.CL_DEVICE_TYPE_ALL;
import static org.jocl.CL.CL_MEM_COPY_HOST_PTR;
import static org.jocl.CL.CL_MEM_READ_ONLY;
import static org.jocl.CL.CL_MEM_READ_WRITE;
import static org.jocl.CL.CL_MEM_WRITE_ONLY;
import static org.jocl.CL.CL_TRUE;
import static org.jocl.CL.clBuildProgram;
import static org.jocl.CL.clCreateBuffer;
import static org.jocl.CL.clCreateCommandQueue;
import static org.jocl.CL.clCreateContext;
import static org.jocl.CL.clCreateKernel;
import static org.jocl.CL.clCreateProgramWithSource;
import static org.jocl.CL.clEnqueueNDRangeKernel;
import static org.jocl.CL.clEnqueueReadBuffer;
import static org.jocl.CL.clGetDeviceIDs;
import static org.jocl.CL.clGetPlatformIDs;
import static org.jocl.CL.clReleaseCommandQueue;
import static org.jocl.CL.clReleaseContext;
import static org.jocl.CL.clReleaseKernel;
import static org.jocl.CL.clReleaseMemObject;
import static org.jocl.CL.clReleaseProgram;
import static org.jocl.CL.clSetKernelArg;
import hpc.exercices.Profiler;
import hpc.exercices.Utils;

import java.io.IOException;
import java.nio.charset.Charset;
import java.util.ArrayList;
import java.util.Arrays;
import java.util.List;

import org.jocl.CL;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_context_properties;
import org.jocl.cl_device_id;
import org.jocl.cl_event;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;
import org.junit.Assert;

/**
 * Exercise 2 - Scan large arrays
 * 
 * @author Hinterleitner Michael, Heily Daniel
 *
 */

public class ScanLargeArrays {

	/*
	 * open cl variables
	 */
	private static cl_context context;

	private static cl_command_queue commandQueue;

	private static cl_program program;

	private static cl_device_id device;
	
	// kernels
	private static cl_kernel kernel;
	private static cl_kernel kernelSmall;
	private static cl_kernel blockSumKernel;

	// array elements
	private static int elements = 1000000;

	// block size
	private static int blockSize = 1024;

	private static long gpuTime;

	public static void main(String[] args) {

		int[] array = Utils.randomArray(elements);

		// ////////////////////////////////////////////////////////////
		// SEQUENTIAL //

		int[] result_seq;

		long startSeqScan = System.currentTimeMillis();
		result_seq = SequentialScan.scan(array);
		long endSeqScan = System.currentTimeMillis();
		System.out.println("Duration for sequential Scan: "
				+ (endSeqScan - startSeqScan) + " ms");

		int[] output = scan(array);
		output = Arrays.copyOfRange(output, 0, elements);
		// System.out.println(Arrays.toString(result_seq));
		// ////////////////////////////////////////////////////////////
		Assert.assertArrayEquals(result_seq, output);

		System.out.println("Scan for " + elements
				+ " elements was performed correctly");

	}

	public static int[] scan(int[] input) {

		setUp();
		createKernels();

		int output[] = new int[input.length];

		// calculate number of required blocks.
		int newSize = input.length;
		double x = Math.log(input.length) / Math.log(2);
		if (Math.abs(x - (int) x) > 0.000000001) {
			newSize = (int) Math.pow(2, ((int) x) + 1);

			int[] newInput = new int[newSize];
			for (int i = 0; i < input.length; i++) {
				newInput[i] = input[i];
			}
			input = newInput;
		}

		int numberOfBlocks = newSize / blockSize;
		
		// debug
		// System.out.println("Number of blocks " + numberOfBlocks);

		if (numberOfBlocks > 1) {
			output = doScanForLargeArray(input, numberOfBlocks);
		} else {
			output = doScanForSmallArray(input);
		}

		shutdown();

		System.out.println("Duration for parallel Scan: "
				+ String.format("%8.3f", gpuTime / 1e6) + " ms");
		return output;
	}

	private static int[] doScanForSmallArray(int[] input) {

		int output[] = new int[input.length];

		Pointer inputPointer = Pointer.to(input);
		Pointer outputPointer = Pointer.to(output);

		// allocate the memory objects for the input- and output data
		cl_mem memObjects[] = new cl_mem[2];
		memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * input.length,
				inputPointer, null);
		memObjects[1] = clCreateBuffer(context, CL_MEM_WRITE_ONLY
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * output.length,
				outputPointer, null);

		// global worksize
		long globalWorkSize[] = new long[] { input.length / 2 };

		// local worksize
		long localWorkSize[] = new long[] { input.length / 2 };

		// setting arguments for kernel
		clSetKernelArg(kernelSmall, 0, Sizeof.cl_mem, Pointer.to(memObjects[0]));
		clSetKernelArg(kernelSmall, 1, Sizeof.cl_mem, Pointer.to(memObjects[1]));
		// temp array in kernel
		// because of ---> variable length array declaration not allowed in
		// OpenCL
		clSetKernelArg(kernelSmall, 2, Sizeof.cl_int * input.length * 2, null);
		clSetKernelArg(kernelSmall, 3, Sizeof.cl_int,
				Pointer.to(new int[] { input.length }));

		cl_event kernel_event = new cl_event();

		// execute the kernel
		clEnqueueNDRangeKernel(commandQueue, kernelSmall, 1, null,
				globalWorkSize, localWorkSize, 0, null, kernel_event);

		// wait for kernel "end" event
		CL.clWaitForEvents(1, new cl_event[] { kernel_event });

		// read the output data
		clEnqueueReadBuffer(commandQueue, memObjects[1], CL_TRUE, 0,
				output.length * Sizeof.cl_int, outputPointer, 0, null, null);

		// release kernel, program, and memory objects
		clReleaseMemObject(memObjects[0]);
		clReleaseMemObject(memObjects[1]);

		gpuTime += Profiler.profile(kernel_event, false);

		return output;
	}

	private static int[] doScanForLargeArray(int[] input, int numberOfBlocks) {

		int[] output = new int[input.length];

		// buffers for the input parts
		cl_mem[] inputBuffers = new cl_mem[numberOfBlocks];
		// buffers for the scans of the single input parts
		cl_mem[] scanBlockBuffers = new cl_mem[numberOfBlocks];
		// buffer which stores sums of each input part
		cl_mem auxiliaryArray;
		// buffer which stores scan of auxiliary array
		cl_mem scanedBlockSumsMem;
		// buffers for the output parts
		cl_mem[] outputBuffer = new cl_mem[numberOfBlocks];

		// arrays for inputs, scannedInput parts, output parts
		List<int[]> inputParts = new ArrayList<int[]>();
		List<int[]> scanBlockBufferParts = new ArrayList<int[]>();
		List<int[]> outputParts = new ArrayList<int[]>();

		// pointers for input parts, output parts and scanned input parts
		// mem_objects
		Pointer[] inputPointers = new Pointer[numberOfBlocks];
		Pointer[] outputPointers = new Pointer[numberOfBlocks];
		Pointer[] scanBlockBuffersPointers = new Pointer[numberOfBlocks];

		int[] auxArray = new int[numberOfBlocks];
		Pointer auxPointer = Pointer.to(auxArray);

		for (int i = 0; i < numberOfBlocks; i++) {
			int[] part = Arrays.copyOfRange(input, i * blockSize, (i + 1)
					* blockSize);
			inputParts.add(part);
			int[] scanBlockBuffersPart = new int[blockSize];
			scanBlockBufferParts.add(scanBlockBuffersPart);
			scanBlockBuffersPointers[i] = Pointer.to(scanBlockBuffersPart);
			int[] outputPart = new int[blockSize];
			Pointer outputPointer = Pointer.to(outputPart);
			outputParts.add(outputPart);
			outputPointers[i] = outputPointer;
		}

		for (int i = 0; i < inputParts.size(); i++) {
			inputPointers[i] = Pointer.to(inputParts.get(i));
		}

		for (int i = 0; i < numberOfBlocks; i++) {

			inputBuffers[i] = clCreateBuffer(context, CL_MEM_READ_ONLY
					| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * blockSize,
					inputPointers[i], null);
			scanBlockBuffers[i] = clCreateBuffer(context, CL_MEM_READ_WRITE
					| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * blockSize,
					scanBlockBuffersPointers[i], null);
			outputBuffer[i] = clCreateBuffer(context, CL_MEM_READ_WRITE,
					Sizeof.cl_int * blockSize, null, null);
		}

		auxiliaryArray = clCreateBuffer(context, CL_MEM_READ_WRITE
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * numberOfBlocks,
				auxPointer, null);

		for (int i = 0; i < numberOfBlocks; i++) {

			// global worksize
			long globalWorkSize[] = new long[] { blockSize / 2 };

			// local worksize
			long localWorkSize[] = new long[] { blockSize / 2 };

			// setting arguments for kernel
			clSetKernelArg(kernel, 0, Sizeof.cl_mem,
					Pointer.to(inputBuffers[i]));
			clSetKernelArg(kernel, 1, Sizeof.cl_mem,
					Pointer.to(scanBlockBuffers[i]));
			// temp array in kernel
			// because of ---> variable length array declaration not allowed in
			// OpenCL

			clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(auxiliaryArray));

			clSetKernelArg(kernel, 3, Sizeof.cl_int * blockSize * 2, null);
			clSetKernelArg(kernel, 4, Sizeof.cl_int,
					Pointer.to(new int[] { blockSize }));
			clSetKernelArg(kernel, 5, Sizeof.cl_int,
					Pointer.to(new int[] { i }));

			cl_event kernel_event = new cl_event();

			// execute the kernel
			clEnqueueNDRangeKernel(commandQueue, kernel, 1, null,
					globalWorkSize, localWorkSize, 0, null, kernel_event);

			// wait for kernel "end" event
			CL.clWaitForEvents(1, new cl_event[] { kernel_event });

			clEnqueueReadBuffer(commandQueue, scanBlockBuffers[i], CL_TRUE, 0,
					blockSize * Sizeof.cl_int, scanBlockBuffersPointers[i], 0,
					null, null);

			gpuTime += Profiler.profile(kernel_event, false);
		}

		clEnqueueReadBuffer(commandQueue, auxiliaryArray, CL_TRUE, 0,
				numberOfBlocks * Sizeof.cl_int, auxPointer, 0, null, null);

		// STEP 2:
		// Scan auxiliary array
		// OUTPUT : scan of auxiliary array in scanendBlockSumsMem

		int[] scanedBlockSums = new int[auxArray.length];

		// check if largeArray scan is necessary or not
		if (auxArray.length / 2 <= blockSize) {
			scanedBlockSums = doScanForSmallArray(auxArray);
		} else {
			scanedBlockSums = doScanForLargeArray(auxArray, auxArray.length
					/ blockSize);
		}

		Pointer scanedBlockSumsPointer = Pointer.to(scanedBlockSums);
		scanedBlockSumsMem = clCreateBuffer(context, CL_MEM_READ_WRITE
				| CL_MEM_COPY_HOST_PTR, Sizeof.cl_int * numberOfBlocks,
				scanedBlockSumsPointer, null);

		// STEP 3:
		// combined scan of input parts with scan of input part sums..
		// das Element der gescannten inputPart Summen mit Index blockNumber
		// wird zu jedem Element des gescannten input parts [blockNumber + 1]
		// hinzugefuegt
		// OUTPUT einzelne output mems + kombination in einem output array
		for (int i = 0; i < numberOfBlocks; i++) {
			long[] globalWorkSize = new long[] { blockSize / 2 };

			// local worksize
			long[] localWorkSize = new long[] { blockSize / 2 };

			// setting arguments for kernel
			clSetKernelArg(blockSumKernel, 0, Sizeof.cl_mem,
					Pointer.to(scanBlockBuffers[i]));
			clSetKernelArg(blockSumKernel, 1, Sizeof.cl_mem,
					Pointer.to(scanedBlockSumsMem));

			clSetKernelArg(blockSumKernel, 2, Sizeof.cl_mem,
					Pointer.to(outputBuffer[i]));

			clSetKernelArg(blockSumKernel, 3, Sizeof.cl_int,
					Pointer.to(new int[] { i }));
			clSetKernelArg(blockSumKernel, 4, Sizeof.cl_int,
					Pointer.to(new int[] { blockSize }));

			cl_event kernel_event = new cl_event();

			// execute the kernel
			clEnqueueNDRangeKernel(commandQueue, blockSumKernel, 1, null,
					globalWorkSize, localWorkSize, 0, null, kernel_event);

			// wait for kernel "end" event
			CL.clWaitForEvents(1, new cl_event[] { kernel_event });

			clEnqueueReadBuffer(commandQueue, outputBuffer[i], CL_TRUE, 0,
					blockSize * Sizeof.cl_int, outputPointers[i], 0, null, null);
		
			System.arraycopy(Arrays.copyOfRange(outputParts.get(i), 0, blockSize),
					0, output, i * blockSize, blockSize);
			
			gpuTime += Profiler.profile(kernel_event, false);
		}

		for (int i = 0; i < numberOfBlocks; i++) {
			clReleaseMemObject(inputBuffers[i]);
			clReleaseMemObject(outputBuffer[i]);
			clReleaseMemObject(scanBlockBuffers[i]);
		}

		clReleaseMemObject(scanedBlockSumsMem);
		clReleaseMemObject(auxiliaryArray);

		return output;
	}

	private static void setUp() {
		final int platformIndex = 0;
		final long deviceType = CL_DEVICE_TYPE_ALL;
		final int deviceIndex = 0;

		// enable exceptions
		CL.setExceptionsEnabled(true);

		// obtain the number of platforms
		int numPlatformsArray[] = new int[1];
		clGetPlatformIDs(0, null, numPlatformsArray);
		int numPlatforms = numPlatformsArray[0];

		// obtain a platform ID
		cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
		clGetPlatformIDs(platforms.length, platforms, null);
		cl_platform_id platform = platforms[platformIndex];

		// initialize the context properties
		cl_context_properties contextProperties = new cl_context_properties();
		contextProperties.addProperty(CL_CONTEXT_PLATFORM, platform);

		// obtain the number of devices for the platform
		int numDevicesArray[] = new int[1];
		clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
		int numDevices = numDevicesArray[0];

		// obtain a device ID
		cl_device_id devices[] = new cl_device_id[numDevices];
		clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
		device = devices[deviceIndex];

		// create a context for the selected device
		context = clCreateContext(contextProperties, 1,
				new cl_device_id[] { device }, null, null, null);

		long properties = 0;
		properties |= CL.CL_QUEUE_PROFILING_ENABLE;
		properties |= CL.CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;

		// create a command-queue for the selected device cl_command_queue
		commandQueue = clCreateCommandQueue(context, device, properties, null);
	}

	private static void createKernels() {

		String programSource = null;
		try {
			programSource = Utils.readFile(
					"kernelCode/exercise2_LargeArrays.cl",
					Charset.defaultCharset());

		} catch (IOException e) {
			System.err.println("kernel source code couldn't be read! Message:"
					+ e.getMessage());
		}

		// create the program from the source code
		program = clCreateProgramWithSource(context, 1,
				new String[] { programSource }, null, null);

		// build the program
		clBuildProgram(program, 0, null, null, null, null);

		// create the kernel
		kernel = clCreateKernel(program, "scan", null);
		kernelSmall = clCreateKernel(program, "scanSmall", null);
		blockSumKernel = clCreateKernel(program, "addBlockSums", null);
	}

	private static void shutdown() {
		clReleaseKernel(kernel);
		clReleaseKernel(kernelSmall);
		clReleaseKernel(blockSumKernel);
		clReleaseProgram(program);
		clReleaseCommandQueue(commandQueue);
		clReleaseContext(context);
	}
}
